#include <Display/Renders/PointRender/point_render_cuda.h>
#include <cuda_block_size_set.h>
#include <Utilities/cuda_utilities.h>
#include <cuda_gl_interop.h>
using namespace std;

namespace PhysLeo {

template<typename T>
PointRenderCuda<T>::PointRenderCuda():PointRenderCpu<T>(),vbo_cuda_(nullptr)
{
}

template<typename T>
PointRenderCuda<T>::~PointRenderCuda()
{
}

template<typename T>
void PointRenderCuda<T>::init(SimData<T> &simData)
{
	PointRenderCpu<T>::init(simData);
    PointRenderCpu<T>::ptr_pos_copy_->allocateGpu();
	cudaGraphicsGLRegisterBuffer(&vbo_cuda_, PointRenderCpu<T>::vbo_, cudaGraphicsMapFlagsNone);
}

template<typename T>
__global__ static void K_PointRenderCuda_transformToDrawData(glm::tvec3<T>* src, glm::vec3* dst, int size)
{
    int pid = threadIdx.x + (blockIdx.x * blockDim.x);
    if (pid >= size) return;

    dst[pid][0] = static_cast<float>(src[pid][0]);
    dst[pid][1] = static_cast<float>(src[pid][1]);
    dst[pid][2] = static_cast<float>(src[pid][2]);
}

template<typename T>
void PointRenderCuda<T>::transformToDrawData()
{
    K_PointRenderCuda_transformToDrawData << <numBlocks(PointRenderCpu<T>::num_points_), BLOCK_SIZE >> > (PointRenderCpu<T>::ptr_pos_->ptrGpu(), PointRenderCpu<T>::ptr_pos_copy_->ptrGpu(), PointRenderCpu<T>::num_points_);
    synchronCheck;

    #ifndef NDEBUG
    this->ptr_pos_copy_->gpu2Cpu();
    #endif 

    PointRenderCpu<T>::exportPointXyz();
}

template<typename T>
void PointRenderCuda<T>::updateDrawData()
{
    void * vbo_device_ptr = nullptr;
	size_t size = 0;
	cudaGraphicsMapResources(1, &vbo_cuda_, nullptr);
	cudaGraphicsResourceGetMappedPointer(&vbo_device_ptr, &size, vbo_cuda_);

	cudaMemcpy(vbo_device_ptr, PointRenderCpu<T>::ptr_pos_copy_->ptrGpu(), PointRenderCpu<T>::num_points_ * sizeof(glm::vec3), cudaMemcpyDeviceToDevice);

	cudaGraphicsUnmapResources(1, &vbo_cuda_, nullptr);
}

template class PointRenderCuda<float>;
template class PointRenderCuda<double>;
}